/*
    Copyright (C) 2009 Giacomo Spigler

    This program is free software; you can redistribute it and/or modify
    it under the terms of the GNU General Public License as published by
    the Free Software Foundation; either version 2 of the License, or
    (at your option) any later version.

    This program is distributed in the hope that it will be useful,
    but WITHOUT ANY WARRANTY; without even the implied warranty of
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
    GNU General Public License for more details.

    You should have received a copy of the GNU General Public License along
    with this program; if not, write to the Free Software Foundation, Inc.,
    51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
*/


#include <stdlib.h>
#include <stdio.h>
#include <time.h>
#include <math.h>

#include <cv.h>

#include <cutil.h>

#include "standard.h"
#include "gabor.h"


texture<unsigned char, 2, cudaReadModeElementType> s2;




__global__ void C2_GPU(int *w_dev, int *h_dev, int nums2, unsigned char *c2) {
  unsigned int p, scale, i;
  int x, y;
  unsigned char v;

  int w[NUMSCALES-1];
  int h[NUMSCALES-1];
  for(scale=0; scale<NUMSCALES-1; scale++) {
    w[scale]=(w_dev[scale]-2*GABOR_W-C1_SIZE)/C1_STEP + 1 -4;
    h[scale]=(h_dev[scale]-2*GABOR_W-C1_SIZE)/C1_STEP + 1 -4;
  }


  __shared__ unsigned char mmax[C2THREADS];

  for(p=blockIdx.x; p<nums2; p+=gridDim.x) {
    unsigned char m=0;

    for(scale=0; scale<NUMSCALES-1; scale++) {
      int sz=w[scale]*h[scale];
      int baseY=h[0]*scale;

      if(w[scale]>=0 && h[scale]>=0) {

        for(i=threadIdx.x; i<sz; i+=blockDim.x) {
          x=i % w[scale];
          y=i / w[scale];

          v=tex2D(s2, x*nums2+p, baseY+y);
          if(v>m) m=v;
        }

      }

    }



    mmax[threadIdx.x]=m;
    __syncthreads();

    if(threadIdx.x==0) {
      #pragma unroll
      for(i=1; i<C2THREADS; i++) {
        if(mmax[i]>m) m=mmax[i];
      }

      c2[p]=m;

    }


  }

}






extern "C" {



void ComputeC2_4(int w_host[NUMSCALES], int h_host[NUMSCALES], int *w_dev, int *h_dev, size_t s2_pitch, unsigned char *s2_dev, int nums2, unsigned char *c2_dev, unsigned char *c2_host,   IplImage *pyr[NUMSCALES]) {
  int w_0=(w_host[0]-2*GABOR_W-C1_SIZE)/C1_STEP + 1 -4;
  int h_0=(h_host[0]-2*GABOR_W-C1_SIZE)/C1_STEP + 1 -4;

  cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>(); //(8, 0, 0, 0, 
  cudaBindTexture2D(0, &s2, s2_dev, &channelDesc, w_0*nums2, h_0*(NUMSCALES-1), s2_pitch);



  C2_GPU<<<C2BLOCKS, C2THREADS>>>(w_dev, h_dev, nums2, c2_dev);
  cudaThreadSynchronize();


  cudaMemcpy(c2_host, c2_dev, nums2*sizeof(unsigned char), cudaMemcpyDeviceToHost);



//TO CHECK
//for(int i=0; i<10; i++) printf("%d\n", c2_host[i]);



  cudaUnbindTexture(s2);

}










}


